/*
 * Copyright (C) 2015 ecl1pse
 *
 * This program is free software; you can redistribute it and/or modify
 * it under the terms of the GNU General Public License as published by
 * the Free Software Foundation; either version 2 of the License, or
 * (at your option) any later version.
 *
 * This program is distributed in the hope that it will be useful,
 * but WITHOUT ANY WARRANTY; without even the implied warranty of
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
 * GNU General Public License for more details.
 *
 * You should have received a copy of the GNU General Public License
 * along with this program; if not, write to the Free Software
 * Foundation, Inc., 59 Temple Place - Suite 330, Boston, MA 02111-1307, USA.
 */

/*
 * This is a little modification of reference implementation of Tiger hash
 * (which you could found here http://www.cs.technion.ac.il/~biham/Reports/Tiger/)
 */
typedef ulong word64;
typedef uchar byte;

#ifdef __ENDIAN_LITTLE__
#define LITTLE_ENDIAN
#undef BIG_ENDIAN
#else
#define BIG_ENDIAN
#undef LITTLE_ENDIAN
#endif

#define TH_BITS 192
#define TH_BYTES 24

#define S_BOX_SIZE 256
#define S_BOX_CNT 4

#define PASSES 3

#define get_byte(unum, i) ( unum >> i*8 & 0xff )

__constant word64 sb_table[];

#define t1 (sb_table)
#define t2 (sb_table + S_BOX_SIZE)
#define t3 (sb_table + S_BOX_SIZE*2)
#define t4 (sb_table + S_BOX_SIZE*3)

#define save_abc \
  aa = a; \
  bb = b; \
  cc = c;
  
#define round(a, b, c, x, mul) \
  c ^= x; \
  a -= t1[get_byte(c, 0)] ^ t2[get_byte(c, 2)] \
    ^ t3[get_byte(c, 4)] ^ t4[get_byte(c, 6)]; \
  b += t4[get_byte(c, 1)] ^ t3[get_byte(c, 3)] \
    ^ t2[get_byte(c, 5)] ^ t1[get_byte(c, 7)]; \
  b *= mul;
  
#define pass(a,b,c,mul) \
  round(a, b, c, x0, mul) \
  round(b, c, a, x1, mul) \
  round(c, a, b, x2, mul) \
  round(a, b, c, x3, mul) \
  round(b, c, a, x4, mul) \
  round(c, a, b, x5, mul) \
  round(a, b, c, x6, mul) \
  round(b, c, a, x7, mul)
  
#define key_schedule \
  x0 -= x7 ^ 0xA5A5A5A5A5A5A5A5UL; \
  x1 ^= x0; \
  x2 += x1; \
  x3 -= x2 ^ (~x1 << 19); \
  x4 ^= x3; \
  x5 += x4; \
  x6 -= x5 ^ (~x4 >> 23); \
  x7 ^= x6; \
  x0 += x7; \
  x1 -= x0 ^ (~x7 << 19); \
  x2 ^= x1; \
  x3 += x2; \
  x4 -= x3 ^ (~x2 >> 23); \
  x5 ^= x4; \
  x6 += x5; \
  x7 -= x6 ^ 0x0123456789ABCDEFUL;
  
#define feedforward \
  a ^= aa; \
  b -= bb; \
  c += cc;
  
#define compress \
  save_abc \
  pass(a, b, c, 5) \
  key_schedule \
  pass(c, a, b, 7) \
  key_schedule \
  pass(b, c, a, 9) \
  for(pass_no = 3; pass_no < PASSES; pass_no++) { \
    key_schedule \
	pass(a,b,c,9) \
	tmpa=a; a=c; c=b; b=tmpa; \
  } \
  feedforward
  
#define tiger_compress_macro(str, state) \
  a = state[0]; \
  b = state[1]; \
  c = state[2]; \
\
  x0=str[0]; x1=str[1]; x2=str[2]; x3=str[3]; \
  x4=str[4]; x5=str[5]; x6=str[6]; x7=str[7]; \
\
  compress; \
\
  state[0] = a; \
  state[1] = b; \
  state[2] = c;

#ifdef BIG_ENDIAN
#define tiger_macro(__adr_spc1, __adr_spc2) \
  word64 i, j; \
  byte temp[64]; \
 \
  word64 a, b, c, tmpa; \
  word64 aa, bb, cc; \
  word64 x0, x1, x2, x3, x4, x5, x6, x7; \
  int pass_no; \
 \
  res[0]=0x0123456789ABCDEFUL; \
  res[1]=0xFEDCBA9876543210UL; \
  res[2]=0xF096A5B4C3B2E187UL; \
 \
  for (i = length; i >= 64; i -= 64) { \
    for (j = 0; j < 64; j++) temp[j^7] = ((__adr_spc1 byte *) str)[j]; \
    tiger_compress_macro(((word64 *) temp), res); \
 \
    str += 8; \
  } \
 \
  for (j = 0; j < i; j++) temp[j^7] = ((__adr_spc1 byte *) str)[j]; \
  temp[j^7] = 0x01; \
  j++; \
 \
  for (; j & 7; j++) temp[j^7] = 0; \
 \
  if (j > 56) { \
    for(; j < 64; j++) temp[j] = 0; \
 \
    tiger_compress_macro(((word64 *)temp), res); \
    j = 0; \
  } \
 \
  for(; j < 56; j++) temp[j] = 0; \
 \
  ((word64 *) (&(temp[56])))[0] = ((word64) length) << 3; \
  tiger_compress_macro(((word64 *) temp), res);

#else
#define tiger_macro(__adr_spc1, __adr_spc2) \
  word64 i, j; \
  byte temp[64]; \
 \
  word64 a, b, c, tmpa; \
  word64 aa, bb, cc; \
  word64 x0, x1, x2, x3, x4, x5, x6, x7; \
  int pass_no; \
 \
  res[0]=0x0123456789ABCDEFUL; \
  res[1]=0xFEDCBA9876543210UL; \
  res[2]=0xF096A5B4C3B2E187UL; \
 \
  for (i = length; i >= 64; i -= 64) { \
    tiger_compress_macro(str, res); \
    str += 8; \
  } \
 \
  for (j = 0; j < i; j++) temp[j] = ((__adr_spc1 byte *) str)[j]; \
 \
  temp[j++] = 0x01; \
  for (; j & 7; j++) temp[j] = 0; \
 \
  if (j > 56) { \
    for(; j < 64; j++) temp[j] = 0; \
 \
    tiger_compress_macro(((word64 *)temp), res); \
    j = 0; \
  } \
 \
  for(; j < 56; j++) temp[j] = 0; \
 \
  ((word64 *) (&(temp[56])))[0] = ((word64) length) << 3; \
  tiger_compress_macro(((word64 *) temp), res);
#endif


/*
 * Computes tiger hash of data.
 * __in str - input message data that needed to be tiger-hashed.
 * __in length - length of data (in bytes).
 * __out res - three 64-bit words of result hash value.
 */
void tiger_pp(word64 *str, word64 length, word64 res[3]) { tiger_macro(__private, __private) }
void tiger_pg(word64 *str, word64 length, __global word64 res[3]) { tiger_macro(__private, __global) }
void tiger_gp(__global word64 *str, word64 length, word64 res[3]) { tiger_macro(__global, __private) }
void tiger_gg(__global word64 *str, word64 length, __global word64 res[3]) { tiger_macro(__global, __global) }

__constant word64 sb_table[S_BOX_CNT*S_BOX_SIZE] = {
    0x02AAB17CF7E90C5EUL   /*    0 */,    0xAC424B03E243A8ECUL   /*    1 */,
    0x72CD5BE30DD5FCD3UL   /*    2 */,    0x6D019B93F6F97F3AUL   /*    3 */,
    0xCD9978FFD21F9193UL   /*    4 */,    0x7573A1C9708029E2UL   /*    5 */,
    0xB164326B922A83C3UL   /*    6 */,    0x46883EEE04915870UL   /*    7 */,
    0xEAACE3057103ECE6UL   /*    8 */,    0xC54169B808A3535CUL   /*    9 */,
    0x4CE754918DDEC47CUL   /*   10 */,    0x0AA2F4DFDC0DF40CUL   /*   11 */,
    0x10B76F18A74DBEFAUL   /*   12 */,    0xC6CCB6235AD1AB6AUL   /*   13 */,
    0x13726121572FE2FFUL   /*   14 */,    0x1A488C6F199D921EUL   /*   15 */,
    0x4BC9F9F4DA0007CAUL   /*   16 */,    0x26F5E6F6E85241C7UL   /*   17 */,
    0x859079DBEA5947B6UL   /*   18 */,    0x4F1885C5C99E8C92UL   /*   19 */,
    0xD78E761EA96F864BUL   /*   20 */,    0x8E36428C52B5C17DUL   /*   21 */,
    0x69CF6827373063C1UL   /*   22 */,    0xB607C93D9BB4C56EUL   /*   23 */,
    0x7D820E760E76B5EAUL   /*   24 */,    0x645C9CC6F07FDC42UL   /*   25 */,
    0xBF38A078243342E0UL   /*   26 */,    0x5F6B343C9D2E7D04UL   /*   27 */,
    0xF2C28AEB600B0EC6UL   /*   28 */,    0x6C0ED85F7254BCACUL   /*   29 */,
    0x71592281A4DB4FE5UL   /*   30 */,    0x1967FA69CE0FED9FUL   /*   31 */,
    0xFD5293F8B96545DBUL   /*   32 */,    0xC879E9D7F2A7600BUL   /*   33 */,
    0x860248920193194EUL   /*   34 */,    0xA4F9533B2D9CC0B3UL   /*   35 */,
    0x9053836C15957613UL   /*   36 */,    0xDB6DCF8AFC357BF1UL   /*   37 */,
    0x18BEEA7A7A370F57UL   /*   38 */,    0x037117CA50B99066UL   /*   39 */,
    0x6AB30A9774424A35UL   /*   40 */,    0xF4E92F02E325249BUL   /*   41 */,
    0x7739DB07061CCAE1UL   /*   42 */,    0xD8F3B49CECA42A05UL   /*   43 */,
    0xBD56BE3F51382F73UL   /*   44 */,    0x45FAED5843B0BB28UL   /*   45 */,
    0x1C813D5C11BF1F83UL   /*   46 */,    0x8AF0E4B6D75FA169UL   /*   47 */,
    0x33EE18A487AD9999UL   /*   48 */,    0x3C26E8EAB1C94410UL   /*   49 */,
    0xB510102BC0A822F9UL   /*   50 */,    0x141EEF310CE6123BUL   /*   51 */,
    0xFC65B90059DDB154UL   /*   52 */,    0xE0158640C5E0E607UL   /*   53 */,
    0x884E079826C3A3CFUL   /*   54 */,    0x930D0D9523C535FDUL   /*   55 */,
    0x35638D754E9A2B00UL   /*   56 */,    0x4085FCCF40469DD5UL   /*   57 */,
    0xC4B17AD28BE23A4CUL   /*   58 */,    0xCAB2F0FC6A3E6A2EUL   /*   59 */,
    0x2860971A6B943FCDUL   /*   60 */,    0x3DDE6EE212E30446UL   /*   61 */,
    0x6222F32AE01765AEUL   /*   62 */,    0x5D550BB5478308FEUL   /*   63 */,
    0xA9EFA98DA0EDA22AUL   /*   64 */,    0xC351A71686C40DA7UL   /*   65 */,
    0x1105586D9C867C84UL   /*   66 */,    0xDCFFEE85FDA22853UL   /*   67 */,
    0xCCFBD0262C5EEF76UL   /*   68 */,    0xBAF294CB8990D201UL   /*   69 */,
    0xE69464F52AFAD975UL   /*   70 */,    0x94B013AFDF133E14UL   /*   71 */,
    0x06A7D1A32823C958UL   /*   72 */,    0x6F95FE5130F61119UL   /*   73 */,
    0xD92AB34E462C06C0UL   /*   74 */,    0xED7BDE33887C71D2UL   /*   75 */,
    0x79746D6E6518393EUL   /*   76 */,    0x5BA419385D713329UL   /*   77 */,
    0x7C1BA6B948A97564UL   /*   78 */,    0x31987C197BFDAC67UL   /*   79 */,
    0xDE6C23C44B053D02UL   /*   80 */,    0x581C49FED002D64DUL   /*   81 */,
    0xDD474D6338261571UL   /*   82 */,    0xAA4546C3E473D062UL   /*   83 */,
    0x928FCE349455F860UL   /*   84 */,    0x48161BBACAAB94D9UL   /*   85 */,
    0x63912430770E6F68UL   /*   86 */,    0x6EC8A5E602C6641CUL   /*   87 */,
    0x87282515337DDD2BUL   /*   88 */,    0x2CDA6B42034B701BUL   /*   89 */,
    0xB03D37C181CB096DUL   /*   90 */,    0xE108438266C71C6FUL   /*   91 */,
    0x2B3180C7EB51B255UL   /*   92 */,    0xDF92B82F96C08BBCUL   /*   93 */,
    0x5C68C8C0A632F3BAUL   /*   94 */,    0x5504CC861C3D0556UL   /*   95 */,
    0xABBFA4E55FB26B8FUL   /*   96 */,    0x41848B0AB3BACEB4UL   /*   97 */,
    0xB334A273AA445D32UL   /*   98 */,    0xBCA696F0A85AD881UL   /*   99 */,
    0x24F6EC65B528D56CUL   /*  100 */,    0x0CE1512E90F4524AUL   /*  101 */,
    0x4E9DD79D5506D35AUL   /*  102 */,    0x258905FAC6CE9779UL   /*  103 */,
    0x2019295B3E109B33UL   /*  104 */,    0xF8A9478B73A054CCUL   /*  105 */,
    0x2924F2F934417EB0UL   /*  106 */,    0x3993357D536D1BC4UL   /*  107 */,
    0x38A81AC21DB6FF8BUL   /*  108 */,    0x47C4FBF17D6016BFUL   /*  109 */,
    0x1E0FAADD7667E3F5UL   /*  110 */,    0x7ABCFF62938BEB96UL   /*  111 */,
    0xA78DAD948FC179C9UL   /*  112 */,    0x8F1F98B72911E50DUL   /*  113 */,
    0x61E48EAE27121A91UL   /*  114 */,    0x4D62F7AD31859808UL   /*  115 */,
    0xECEBA345EF5CEAEBUL   /*  116 */,    0xF5CEB25EBC9684CEUL   /*  117 */,
    0xF633E20CB7F76221UL   /*  118 */,    0xA32CDF06AB8293E4UL   /*  119 */,
    0x985A202CA5EE2CA4UL   /*  120 */,    0xCF0B8447CC8A8FB1UL   /*  121 */,
    0x9F765244979859A3UL   /*  122 */,    0xA8D516B1A1240017UL   /*  123 */,
    0x0BD7BA3EBB5DC726UL   /*  124 */,    0xE54BCA55B86ADB39UL   /*  125 */,
    0x1D7A3AFD6C478063UL   /*  126 */,    0x519EC608E7669EDDUL   /*  127 */,
    0x0E5715A2D149AA23UL   /*  128 */,    0x177D4571848FF194UL   /*  129 */,
    0xEEB55F3241014C22UL   /*  130 */,    0x0F5E5CA13A6E2EC2UL   /*  131 */,
    0x8029927B75F5C361UL   /*  132 */,    0xAD139FABC3D6E436UL   /*  133 */,
    0x0D5DF1A94CCF402FUL   /*  134 */,    0x3E8BD948BEA5DFC8UL   /*  135 */,
    0xA5A0D357BD3FF77EUL   /*  136 */,    0xA2D12E251F74F645UL   /*  137 */,
    0x66FD9E525E81A082UL   /*  138 */,    0x2E0C90CE7F687A49UL   /*  139 */,
    0xC2E8BCBEBA973BC5UL   /*  140 */,    0x000001BCE509745FUL   /*  141 */,
    0x423777BBE6DAB3D6UL   /*  142 */,    0xD1661C7EAEF06EB5UL   /*  143 */,
    0xA1781F354DAACFD8UL   /*  144 */,    0x2D11284A2B16AFFCUL   /*  145 */,
    0xF1FC4F67FA891D1FUL   /*  146 */,    0x73ECC25DCB920ADAUL   /*  147 */,
    0xAE610C22C2A12651UL   /*  148 */,    0x96E0A810D356B78AUL   /*  149 */,
    0x5A9A381F2FE7870FUL   /*  150 */,    0xD5AD62EDE94E5530UL   /*  151 */,
    0xD225E5E8368D1427UL   /*  152 */,    0x65977B70C7AF4631UL   /*  153 */,
    0x99F889B2DE39D74FUL   /*  154 */,    0x233F30BF54E1D143UL   /*  155 */,
    0x9A9675D3D9A63C97UL   /*  156 */,    0x5470554FF334F9A8UL   /*  157 */,
    0x166ACB744A4F5688UL   /*  158 */,    0x70C74CAAB2E4AEADUL   /*  159 */,
    0xF0D091646F294D12UL   /*  160 */,    0x57B82A89684031D1UL   /*  161 */,
    0xEFD95A5A61BE0B6BUL   /*  162 */,    0x2FBD12E969F2F29AUL   /*  163 */,
    0x9BD37013FEFF9FE8UL   /*  164 */,    0x3F9B0404D6085A06UL   /*  165 */,
    0x4940C1F3166CFE15UL   /*  166 */,    0x09542C4DCDF3DEFBUL   /*  167 */,
    0xB4C5218385CD5CE3UL   /*  168 */,    0xC935B7DC4462A641UL   /*  169 */,
    0x3417F8A68ED3B63FUL   /*  170 */,    0xB80959295B215B40UL   /*  171 */,
    0xF99CDAEF3B8C8572UL   /*  172 */,    0x018C0614F8FCB95DUL   /*  173 */,
    0x1B14ACCD1A3ACDF3UL   /*  174 */,    0x84D471F200BB732DUL   /*  175 */,
    0xC1A3110E95E8DA16UL   /*  176 */,    0x430A7220BF1A82B8UL   /*  177 */,
    0xB77E090D39DF210EUL   /*  178 */,    0x5EF4BD9F3CD05E9DUL   /*  179 */,
    0x9D4FF6DA7E57A444UL   /*  180 */,    0xDA1D60E183D4A5F8UL   /*  181 */,
    0xB287C38417998E47UL   /*  182 */,    0xFE3EDC121BB31886UL   /*  183 */,
    0xC7FE3CCC980CCBEFUL   /*  184 */,    0xE46FB590189BFD03UL   /*  185 */,
    0x3732FD469A4C57DCUL   /*  186 */,    0x7EF700A07CF1AD65UL   /*  187 */,
    0x59C64468A31D8859UL   /*  188 */,    0x762FB0B4D45B61F6UL   /*  189 */,
    0x155BAED099047718UL   /*  190 */,    0x68755E4C3D50BAA6UL   /*  191 */,
    0xE9214E7F22D8B4DFUL   /*  192 */,    0x2ADDBF532EAC95F4UL   /*  193 */,
    0x32AE3909B4BD0109UL   /*  194 */,    0x834DF537B08E3450UL   /*  195 */,
    0xFA209DA84220728DUL   /*  196 */,    0x9E691D9B9EFE23F7UL   /*  197 */,
    0x0446D288C4AE8D7FUL   /*  198 */,    0x7B4CC524E169785BUL   /*  199 */,
    0x21D87F0135CA1385UL   /*  200 */,    0xCEBB400F137B8AA5UL   /*  201 */,
    0x272E2B66580796BEUL   /*  202 */,    0x3612264125C2B0DEUL   /*  203 */,
    0x057702BDAD1EFBB2UL   /*  204 */,    0xD4BABB8EACF84BE9UL   /*  205 */,
    0x91583139641BC67BUL   /*  206 */,    0x8BDC2DE08036E024UL   /*  207 */,
    0x603C8156F49F68EDUL   /*  208 */,    0xF7D236F7DBEF5111UL   /*  209 */,
    0x9727C4598AD21E80UL   /*  210 */,    0xA08A0896670A5FD7UL   /*  211 */,
    0xCB4A8F4309EBA9CBUL   /*  212 */,    0x81AF564B0F7036A1UL   /*  213 */,
    0xC0B99AA778199ABDUL   /*  214 */,    0x959F1EC83FC8E952UL   /*  215 */,
    0x8C505077794A81B9UL   /*  216 */,    0x3ACAAF8F056338F0UL   /*  217 */,
    0x07B43F50627A6778UL   /*  218 */,    0x4A44AB49F5ECCC77UL   /*  219 */,
    0x3BC3D6E4B679EE98UL   /*  220 */,    0x9CC0D4D1CF14108CUL   /*  221 */,
    0x4406C00B206BC8A0UL   /*  222 */,    0x82A18854C8D72D89UL   /*  223 */,
    0x67E366B35C3C432CUL   /*  224 */,    0xB923DD61102B37F2UL   /*  225 */,
    0x56AB2779D884271DUL   /*  226 */,    0xBE83E1B0FF1525AFUL   /*  227 */,
    0xFB7C65D4217E49A9UL   /*  228 */,    0x6BDBE0E76D48E7D4UL   /*  229 */,
    0x08DF828745D9179EUL   /*  230 */,    0x22EA6A9ADD53BD34UL   /*  231 */,
    0xE36E141C5622200AUL   /*  232 */,    0x7F805D1B8CB750EEUL   /*  233 */,
    0xAFE5C7A59F58E837UL   /*  234 */,    0xE27F996A4FB1C23CUL   /*  235 */,
    0xD3867DFB0775F0D0UL   /*  236 */,    0xD0E673DE6E88891AUL   /*  237 */,
    0x123AEB9EAFB86C25UL   /*  238 */,    0x30F1D5D5C145B895UL   /*  239 */,
    0xBB434A2DEE7269E7UL   /*  240 */,    0x78CB67ECF931FA38UL   /*  241 */,
    0xF33B0372323BBF9CUL   /*  242 */,    0x52D66336FB279C74UL   /*  243 */,
    0x505F33AC0AFB4EAAUL   /*  244 */,    0xE8A5CD99A2CCE187UL   /*  245 */,
    0x534974801E2D30BBUL   /*  246 */,    0x8D2D5711D5876D90UL   /*  247 */,
    0x1F1A412891BC038EUL   /*  248 */,    0xD6E2E71D82E56648UL   /*  249 */,
    0x74036C3A497732B7UL   /*  250 */,    0x89B67ED96361F5ABUL   /*  251 */,
    0xFFED95D8F1EA02A2UL   /*  252 */,    0xE72B3BD61464D43DUL   /*  253 */,
    0xA6300F170BDC4820UL   /*  254 */,    0xEBC18760ED78A77AUL   /*  255 */,
    0xE6A6BE5A05A12138UL   /*  256 */,    0xB5A122A5B4F87C98UL   /*  257 */,
    0x563C6089140B6990UL   /*  258 */,    0x4C46CB2E391F5DD5UL   /*  259 */,
    0xD932ADDBC9B79434UL   /*  260 */,    0x08EA70E42015AFF5UL   /*  261 */,
    0xD765A6673E478CF1UL   /*  262 */,    0xC4FB757EAB278D99UL   /*  263 */,
    0xDF11C6862D6E0692UL   /*  264 */,    0xDDEB84F10D7F3B16UL   /*  265 */,
    0x6F2EF604A665EA04UL   /*  266 */,    0x4A8E0F0FF0E0DFB3UL   /*  267 */,
    0xA5EDEEF83DBCBA51UL   /*  268 */,    0xFC4F0A2A0EA4371EUL   /*  269 */,
    0xE83E1DA85CB38429UL   /*  270 */,    0xDC8FF882BA1B1CE2UL   /*  271 */,
    0xCD45505E8353E80DUL   /*  272 */,    0x18D19A00D4DB0717UL   /*  273 */,
    0x34A0CFEDA5F38101UL   /*  274 */,    0x0BE77E518887CAF2UL   /*  275 */,
    0x1E341438B3C45136UL   /*  276 */,    0xE05797F49089CCF9UL   /*  277 */,
    0xFFD23F9DF2591D14UL   /*  278 */,    0x543DDA228595C5CDUL   /*  279 */,
    0x661F81FD99052A33UL   /*  280 */,    0x8736E641DB0F7B76UL   /*  281 */,
    0x15227725418E5307UL   /*  282 */,    0xE25F7F46162EB2FAUL   /*  283 */,
    0x48A8B2126C13D9FEUL   /*  284 */,    0xAFDC541792E76EEAUL   /*  285 */,
    0x03D912BFC6D1898FUL   /*  286 */,    0x31B1AAFA1B83F51BUL   /*  287 */,
    0xF1AC2796E42AB7D9UL   /*  288 */,    0x40A3A7D7FCD2EBACUL   /*  289 */,
    0x1056136D0AFBBCC5UL   /*  290 */,    0x7889E1DD9A6D0C85UL   /*  291 */,
    0xD33525782A7974AAUL   /*  292 */,    0xA7E25D09078AC09BUL   /*  293 */,
    0xBD4138B3EAC6EDD0UL   /*  294 */,    0x920ABFBE71EB9E70UL   /*  295 */,
    0xA2A5D0F54FC2625CUL   /*  296 */,    0xC054E36B0B1290A3UL   /*  297 */,
    0xF6DD59FF62FE932BUL   /*  298 */,    0x3537354511A8AC7DUL   /*  299 */,
    0xCA845E9172FADCD4UL   /*  300 */,    0x84F82B60329D20DCUL   /*  301 */,
    0x79C62CE1CD672F18UL   /*  302 */,    0x8B09A2ADD124642CUL   /*  303 */,
    0xD0C1E96A19D9E726UL   /*  304 */,    0x5A786A9B4BA9500CUL   /*  305 */,
    0x0E020336634C43F3UL   /*  306 */,    0xC17B474AEB66D822UL   /*  307 */,
    0x6A731AE3EC9BAAC2UL   /*  308 */,    0x8226667AE0840258UL   /*  309 */,
    0x67D4567691CAECA5UL   /*  310 */,    0x1D94155C4875ADB5UL   /*  311 */,
    0x6D00FD985B813FDFUL   /*  312 */,    0x51286EFCB774CD06UL   /*  313 */,
    0x5E8834471FA744AFUL   /*  314 */,    0xF72CA0AEE761AE2EUL   /*  315 */,
    0xBE40E4CDAEE8E09AUL   /*  316 */,    0xE9970BBB5118F665UL   /*  317 */,
    0x726E4BEB33DF1964UL   /*  318 */,    0x703B000729199762UL   /*  319 */,
    0x4631D816F5EF30A7UL   /*  320 */,    0xB880B5B51504A6BEUL   /*  321 */,
    0x641793C37ED84B6CUL   /*  322 */,    0x7B21ED77F6E97D96UL   /*  323 */,
    0x776306312EF96B73UL   /*  324 */,    0xAE528948E86FF3F4UL   /*  325 */,
    0x53DBD7F286A3F8F8UL   /*  326 */,    0x16CADCE74CFC1063UL   /*  327 */,
    0x005C19BDFA52C6DDUL   /*  328 */,    0x68868F5D64D46AD3UL   /*  329 */,
    0x3A9D512CCF1E186AUL   /*  330 */,    0x367E62C2385660AEUL   /*  331 */,
    0xE359E7EA77DCB1D7UL   /*  332 */,    0x526C0773749ABE6EUL   /*  333 */,
    0x735AE5F9D09F734BUL   /*  334 */,    0x493FC7CC8A558BA8UL   /*  335 */,
    0xB0B9C1533041AB45UL   /*  336 */,    0x321958BA470A59BDUL   /*  337 */,
    0x852DB00B5F46C393UL   /*  338 */,    0x91209B2BD336B0E5UL   /*  339 */,
    0x6E604F7D659EF19FUL   /*  340 */,    0xB99A8AE2782CCB24UL   /*  341 */,
    0xCCF52AB6C814C4C7UL   /*  342 */,    0x4727D9AFBE11727BUL   /*  343 */,
    0x7E950D0C0121B34DUL   /*  344 */,    0x756F435670AD471FUL   /*  345 */,
    0xF5ADD442615A6849UL   /*  346 */,    0x4E87E09980B9957AUL   /*  347 */,
    0x2ACFA1DF50AEE355UL   /*  348 */,    0xD898263AFD2FD556UL   /*  349 */,
    0xC8F4924DD80C8FD6UL   /*  350 */,    0xCF99CA3D754A173AUL   /*  351 */,
    0xFE477BACAF91BF3CUL   /*  352 */,    0xED5371F6D690C12DUL   /*  353 */,
    0x831A5C285E687094UL   /*  354 */,    0xC5D3C90A3708A0A4UL   /*  355 */,
    0x0F7F903717D06580UL   /*  356 */,    0x19F9BB13B8FDF27FUL   /*  357 */,
    0xB1BD6F1B4D502843UL   /*  358 */,    0x1C761BA38FFF4012UL   /*  359 */,
    0x0D1530C4E2E21F3BUL   /*  360 */,    0x8943CE69A7372C8AUL   /*  361 */,
    0xE5184E11FEB5CE66UL   /*  362 */,    0x618BDB80BD736621UL   /*  363 */,
    0x7D29BAD68B574D0BUL   /*  364 */,    0x81BB613E25E6FE5BUL   /*  365 */,
    0x071C9C10BC07913FUL   /*  366 */,    0xC7BEEB7909AC2D97UL   /*  367 */,
    0xC3E58D353BC5D757UL   /*  368 */,    0xEB017892F38F61E8UL   /*  369 */,
    0xD4EFFB9C9B1CC21AUL   /*  370 */,    0x99727D26F494F7ABUL   /*  371 */,
    0xA3E063A2956B3E03UL   /*  372 */,    0x9D4A8B9A4AA09C30UL   /*  373 */,
    0x3F6AB7D500090FB4UL   /*  374 */,    0x9CC0F2A057268AC0UL   /*  375 */,
    0x3DEE9D2DEDBF42D1UL   /*  376 */,    0x330F49C87960A972UL   /*  377 */,
    0xC6B2720287421B41UL   /*  378 */,    0x0AC59EC07C00369CUL   /*  379 */,
    0xEF4EAC49CB353425UL   /*  380 */,    0xF450244EEF0129D8UL   /*  381 */,
    0x8ACC46E5CAF4DEB6UL   /*  382 */,    0x2FFEAB63989263F7UL   /*  383 */,
    0x8F7CB9FE5D7A4578UL   /*  384 */,    0x5BD8F7644E634635UL   /*  385 */,
    0x427A7315BF2DC900UL   /*  386 */,    0x17D0C4AA2125261CUL   /*  387 */,
    0x3992486C93518E50UL   /*  388 */,    0xB4CBFEE0A2D7D4C3UL   /*  389 */,
    0x7C75D6202C5DDD8DUL   /*  390 */,    0xDBC295D8E35B6C61UL   /*  391 */,
    0x60B369D302032B19UL   /*  392 */,    0xCE42685FDCE44132UL   /*  393 */,
    0x06F3DDB9DDF65610UL   /*  394 */,    0x8EA4D21DB5E148F0UL   /*  395 */,
    0x20B0FCE62FCD496FUL   /*  396 */,    0x2C1B912358B0EE31UL   /*  397 */,
    0xB28317B818F5A308UL   /*  398 */,    0xA89C1E189CA6D2CFUL   /*  399 */,
    0x0C6B18576AAADBC8UL   /*  400 */,    0xB65DEAA91299FAE3UL   /*  401 */,
    0xFB2B794B7F1027E7UL   /*  402 */,    0x04E4317F443B5BEBUL   /*  403 */,
    0x4B852D325939D0A6UL   /*  404 */,    0xD5AE6BEEFB207FFCUL   /*  405 */,
    0x309682B281C7D374UL   /*  406 */,    0xBAE309A194C3B475UL   /*  407 */,
    0x8CC3F97B13B49F05UL   /*  408 */,    0x98A9422FF8293967UL   /*  409 */,
    0x244B16B01076FF7CUL   /*  410 */,    0xF8BF571C663D67EEUL   /*  411 */,
    0x1F0D6758EEE30DA1UL   /*  412 */,    0xC9B611D97ADEB9B7UL   /*  413 */,
    0xB7AFD5887B6C57A2UL   /*  414 */,    0x6290AE846B984FE1UL   /*  415 */,
    0x94DF4CDEACC1A5FDUL   /*  416 */,    0x058A5BD1C5483AFFUL   /*  417 */,
    0x63166CC142BA3C37UL   /*  418 */,    0x8DB8526EB2F76F40UL   /*  419 */,
    0xE10880036F0D6D4EUL   /*  420 */,    0x9E0523C9971D311DUL   /*  421 */,
    0x45EC2824CC7CD691UL   /*  422 */,    0x575B8359E62382C9UL   /*  423 */,
    0xFA9E400DC4889995UL   /*  424 */,    0xD1823ECB45721568UL   /*  425 */,
    0xDAFD983B8206082FUL   /*  426 */,    0xAA7D29082386A8CBUL   /*  427 */,
    0x269FCD4403B87588UL   /*  428 */,    0x1B91F5F728BDD1E0UL   /*  429 */,
    0xE4669F39040201F6UL   /*  430 */,    0x7A1D7C218CF04ADEUL   /*  431 */,
    0x65623C29D79CE5CEUL   /*  432 */,    0x2368449096C00BB1UL   /*  433 */,
    0xAB9BF1879DA503BAUL   /*  434 */,    0xBC23ECB1A458058EUL   /*  435 */,
    0x9A58DF01BB401ECCUL   /*  436 */,    0xA070E868A85F143DUL   /*  437 */,
    0x4FF188307DF2239EUL   /*  438 */,    0x14D565B41A641183UL   /*  439 */,
    0xEE13337452701602UL   /*  440 */,    0x950E3DCF3F285E09UL   /*  441 */,
    0x59930254B9C80953UL   /*  442 */,    0x3BF299408930DA6DUL   /*  443 */,
    0xA955943F53691387UL   /*  444 */,    0xA15EDECAA9CB8784UL   /*  445 */,
    0x29142127352BE9A0UL   /*  446 */,    0x76F0371FFF4E7AFBUL   /*  447 */,
    0x0239F450274F2228UL   /*  448 */,    0xBB073AF01D5E868BUL   /*  449 */,
    0xBFC80571C10E96C1UL   /*  450 */,    0xD267088568222E23UL   /*  451 */,
    0x9671A3D48E80B5B0UL   /*  452 */,    0x55B5D38AE193BB81UL   /*  453 */,
    0x693AE2D0A18B04B8UL   /*  454 */,    0x5C48B4ECADD5335FUL   /*  455 */,
    0xFD743B194916A1CAUL   /*  456 */,    0x2577018134BE98C4UL   /*  457 */,
    0xE77987E83C54A4ADUL   /*  458 */,    0x28E11014DA33E1B9UL   /*  459 */,
    0x270CC59E226AA213UL   /*  460 */,    0x71495F756D1A5F60UL   /*  461 */,
    0x9BE853FB60AFEF77UL   /*  462 */,    0xADC786A7F7443DBFUL   /*  463 */,
    0x0904456173B29A82UL   /*  464 */,    0x58BC7A66C232BD5EUL   /*  465 */,
    0xF306558C673AC8B2UL   /*  466 */,    0x41F639C6B6C9772AUL   /*  467 */,
    0x216DEFE99FDA35DAUL   /*  468 */,    0x11640CC71C7BE615UL   /*  469 */,
    0x93C43694565C5527UL   /*  470 */,    0xEA038E6246777839UL   /*  471 */,
    0xF9ABF3CE5A3E2469UL   /*  472 */,    0x741E768D0FD312D2UL   /*  473 */,
    0x0144B883CED652C6UL   /*  474 */,    0xC20B5A5BA33F8552UL   /*  475 */,
    0x1AE69633C3435A9DUL   /*  476 */,    0x97A28CA4088CFDECUL   /*  477 */,
    0x8824A43C1E96F420UL   /*  478 */,    0x37612FA66EEEA746UL   /*  479 */,
    0x6B4CB165F9CF0E5AUL   /*  480 */,    0x43AA1C06A0ABFB4AUL   /*  481 */,
    0x7F4DC26FF162796BUL   /*  482 */,    0x6CBACC8E54ED9B0FUL   /*  483 */,
    0xA6B7FFEFD2BB253EUL   /*  484 */,    0x2E25BC95B0A29D4FUL   /*  485 */,
    0x86D6A58BDEF1388CUL   /*  486 */,    0xDED74AC576B6F054UL   /*  487 */,
    0x8030BDBC2B45805DUL   /*  488 */,    0x3C81AF70E94D9289UL   /*  489 */,
    0x3EFF6DDA9E3100DBUL   /*  490 */,    0xB38DC39FDFCC8847UL   /*  491 */,
    0x123885528D17B87EUL   /*  492 */,    0xF2DA0ED240B1B642UL   /*  493 */,
    0x44CEFADCD54BF9A9UL   /*  494 */,    0x1312200E433C7EE6UL   /*  495 */,
    0x9FFCC84F3A78C748UL   /*  496 */,    0xF0CD1F72248576BBUL   /*  497 */,
    0xEC6974053638CFE4UL   /*  498 */,    0x2BA7B67C0CEC4E4CUL   /*  499 */,
    0xAC2F4DF3E5CE32EDUL   /*  500 */,    0xCB33D14326EA4C11UL   /*  501 */,
    0xA4E9044CC77E58BCUL   /*  502 */,    0x5F513293D934FCEFUL   /*  503 */,
    0x5DC9645506E55444UL   /*  504 */,    0x50DE418F317DE40AUL   /*  505 */,
    0x388CB31A69DDE259UL   /*  506 */,    0x2DB4A83455820A86UL   /*  507 */,
    0x9010A91E84711AE9UL   /*  508 */,    0x4DF7F0B7B1498371UL   /*  509 */,
    0xD62A2EABC0977179UL   /*  510 */,    0x22FAC097AA8D5C0EUL   /*  511 */,
    0xF49FCC2FF1DAF39BUL   /*  512 */,    0x487FD5C66FF29281UL   /*  513 */,
    0xE8A30667FCDCA83FUL   /*  514 */,    0x2C9B4BE3D2FCCE63UL   /*  515 */,
    0xDA3FF74B93FBBBC2UL   /*  516 */,    0x2FA165D2FE70BA66UL   /*  517 */,
    0xA103E279970E93D4UL   /*  518 */,    0xBECDEC77B0E45E71UL   /*  519 */,
    0xCFB41E723985E497UL   /*  520 */,    0xB70AAA025EF75017UL   /*  521 */,
    0xD42309F03840B8E0UL   /*  522 */,    0x8EFC1AD035898579UL   /*  523 */,
    0x96C6920BE2B2ABC5UL   /*  524 */,    0x66AF4163375A9172UL   /*  525 */,
    0x2174ABDCCA7127FBUL   /*  526 */,    0xB33CCEA64A72FF41UL   /*  527 */,
    0xF04A4933083066A5UL   /*  528 */,    0x8D970ACDD7289AF5UL   /*  529 */,
    0x8F96E8E031C8C25EUL   /*  530 */,    0xF3FEC02276875D47UL   /*  531 */,
    0xEC7BF310056190DDUL   /*  532 */,    0xF5ADB0AEBB0F1491UL   /*  533 */,
    0x9B50F8850FD58892UL   /*  534 */,    0x4975488358B74DE8UL   /*  535 */,
    0xA3354FF691531C61UL   /*  536 */,    0x0702BBE481D2C6EEUL   /*  537 */,
    0x89FB24057DEDED98UL   /*  538 */,    0xAC3075138596E902UL   /*  539 */,
    0x1D2D3580172772EDUL   /*  540 */,    0xEB738FC28E6BC30DUL   /*  541 */,
    0x5854EF8F63044326UL   /*  542 */,    0x9E5C52325ADD3BBEUL   /*  543 */,
    0x90AA53CF325C4623UL   /*  544 */,    0xC1D24D51349DD067UL   /*  545 */,
    0x2051CFEEA69EA624UL   /*  546 */,    0x13220F0A862E7E4FUL   /*  547 */,
    0xCE39399404E04864UL   /*  548 */,    0xD9C42CA47086FCB7UL   /*  549 */,
    0x685AD2238A03E7CCUL   /*  550 */,    0x066484B2AB2FF1DBUL   /*  551 */,
    0xFE9D5D70EFBF79ECUL   /*  552 */,    0x5B13B9DD9C481854UL   /*  553 */,
    0x15F0D475ED1509ADUL   /*  554 */,    0x0BEBCD060EC79851UL   /*  555 */,
    0xD58C6791183AB7F8UL   /*  556 */,    0xD1187C5052F3EEE4UL   /*  557 */,
    0xC95D1192E54E82FFUL   /*  558 */,    0x86EEA14CB9AC6CA2UL   /*  559 */,
    0x3485BEB153677D5DUL   /*  560 */,    0xDD191D781F8C492AUL   /*  561 */,
    0xF60866BAA784EBF9UL   /*  562 */,    0x518F643BA2D08C74UL   /*  563 */,
    0x8852E956E1087C22UL   /*  564 */,    0xA768CB8DC410AE8DUL   /*  565 */,
    0x38047726BFEC8E1AUL   /*  566 */,    0xA67738B4CD3B45AAUL   /*  567 */,
    0xAD16691CEC0DDE19UL   /*  568 */,    0xC6D4319380462E07UL   /*  569 */,
    0xC5A5876D0BA61938UL   /*  570 */,    0x16B9FA1FA58FD840UL   /*  571 */,
    0x188AB1173CA74F18UL   /*  572 */,    0xABDA2F98C99C021FUL   /*  573 */,
    0x3E0580AB134AE816UL   /*  574 */,    0x5F3B05B773645ABBUL   /*  575 */,
    0x2501A2BE5575F2F6UL   /*  576 */,    0x1B2F74004E7E8BA9UL   /*  577 */,
    0x1CD7580371E8D953UL   /*  578 */,    0x7F6ED89562764E30UL   /*  579 */,
    0xB15926FF596F003DUL   /*  580 */,    0x9F65293DA8C5D6B9UL   /*  581 */,
    0x6ECEF04DD690F84CUL   /*  582 */,    0x4782275FFF33AF88UL   /*  583 */,
    0xE41433083F820801UL   /*  584 */,    0xFD0DFE409A1AF9B5UL   /*  585 */,
    0x4325A3342CDB396BUL   /*  586 */,    0x8AE77E62B301B252UL   /*  587 */,
    0xC36F9E9F6655615AUL   /*  588 */,    0x85455A2D92D32C09UL   /*  589 */,
    0xF2C7DEA949477485UL   /*  590 */,    0x63CFB4C133A39EBAUL   /*  591 */,
    0x83B040CC6EBC5462UL   /*  592 */,    0x3B9454C8FDB326B0UL   /*  593 */,
    0x56F56A9E87FFD78CUL   /*  594 */,    0x2DC2940D99F42BC6UL   /*  595 */,
    0x98F7DF096B096E2DUL   /*  596 */,    0x19A6E01E3AD852BFUL   /*  597 */,
    0x42A99CCBDBD4B40BUL   /*  598 */,    0xA59998AF45E9C559UL   /*  599 */,
    0x366295E807D93186UL   /*  600 */,    0x6B48181BFAA1F773UL   /*  601 */,
    0x1FEC57E2157A0A1DUL   /*  602 */,    0x4667446AF6201AD5UL   /*  603 */,
    0xE615EBCACFB0F075UL   /*  604 */,    0xB8F31F4F68290778UL   /*  605 */,
    0x22713ED6CE22D11EUL   /*  606 */,    0x3057C1A72EC3C93BUL   /*  607 */,
    0xCB46ACC37C3F1F2FUL   /*  608 */,    0xDBB893FD02AAF50EUL   /*  609 */,
    0x331FD92E600B9FCFUL   /*  610 */,    0xA498F96148EA3AD6UL   /*  611 */,
    0xA8D8426E8B6A83EAUL   /*  612 */,    0xA089B274B7735CDCUL   /*  613 */,
    0x87F6B3731E524A11UL   /*  614 */,    0x118808E5CBC96749UL   /*  615 */,
    0x9906E4C7B19BD394UL   /*  616 */,    0xAFED7F7E9B24A20CUL   /*  617 */,
    0x6509EADEEB3644A7UL   /*  618 */,    0x6C1EF1D3E8EF0EDEUL   /*  619 */,
    0xB9C97D43E9798FB4UL   /*  620 */,    0xA2F2D784740C28A3UL   /*  621 */,
    0x7B8496476197566FUL   /*  622 */,    0x7A5BE3E6B65F069DUL   /*  623 */,
    0xF96330ED78BE6F10UL   /*  624 */,    0xEEE60DE77A076A15UL   /*  625 */,
    0x2B4BEE4AA08B9BD0UL   /*  626 */,    0x6A56A63EC7B8894EUL   /*  627 */,
    0x02121359BA34FEF4UL   /*  628 */,    0x4CBF99F8283703FCUL   /*  629 */,
    0x398071350CAF30C8UL   /*  630 */,    0xD0A77A89F017687AUL   /*  631 */,
    0xF1C1A9EB9E423569UL   /*  632 */,    0x8C7976282DEE8199UL   /*  633 */,
    0x5D1737A5DD1F7ABDUL   /*  634 */,    0x4F53433C09A9FA80UL   /*  635 */,
    0xFA8B0C53DF7CA1D9UL   /*  636 */,    0x3FD9DCBC886CCB77UL   /*  637 */,
    0xC040917CA91B4720UL   /*  638 */,    0x7DD00142F9D1DCDFUL   /*  639 */,
    0x8476FC1D4F387B58UL   /*  640 */,    0x23F8E7C5F3316503UL   /*  641 */,
    0x032A2244E7E37339UL   /*  642 */,    0x5C87A5D750F5A74BUL   /*  643 */,
    0x082B4CC43698992EUL   /*  644 */,    0xDF917BECB858F63CUL   /*  645 */,
    0x3270B8FC5BF86DDAUL   /*  646 */,    0x10AE72BB29B5DD76UL   /*  647 */,
    0x576AC94E7700362BUL   /*  648 */,    0x1AD112DAC61EFB8FUL   /*  649 */,
    0x691BC30EC5FAA427UL   /*  650 */,    0xFF246311CC327143UL   /*  651 */,
    0x3142368E30E53206UL   /*  652 */,    0x71380E31E02CA396UL   /*  653 */,
    0x958D5C960AAD76F1UL   /*  654 */,    0xF8D6F430C16DA536UL   /*  655 */,
    0xC8FFD13F1BE7E1D2UL   /*  656 */,    0x7578AE66004DDBE1UL   /*  657 */,
    0x05833F01067BE646UL   /*  658 */,    0xBB34B5AD3BFE586DUL   /*  659 */,
    0x095F34C9A12B97F0UL   /*  660 */,    0x247AB64525D60CA8UL   /*  661 */,
    0xDCDBC6F3017477D1UL   /*  662 */,    0x4A2E14D4DECAD24DUL   /*  663 */,
    0xBDB5E6D9BE0A1EEBUL   /*  664 */,    0x2A7E70F7794301ABUL   /*  665 */,
    0xDEF42D8A270540FDUL   /*  666 */,    0x01078EC0A34C22C1UL   /*  667 */,
    0xE5DE511AF4C16387UL   /*  668 */,    0x7EBB3A52BD9A330AUL   /*  669 */,
    0x77697857AA7D6435UL   /*  670 */,    0x004E831603AE4C32UL   /*  671 */,
    0xE7A21020AD78E312UL   /*  672 */,    0x9D41A70C6AB420F2UL   /*  673 */,
    0x28E06C18EA1141E6UL   /*  674 */,    0xD2B28CBD984F6B28UL   /*  675 */,
    0x26B75F6C446E9D83UL   /*  676 */,    0xBA47568C4D418D7FUL   /*  677 */,
    0xD80BADBFE6183D8EUL   /*  678 */,    0x0E206D7F5F166044UL   /*  679 */,
    0xE258A43911CBCA3EUL   /*  680 */,    0x723A1746B21DC0BCUL   /*  681 */,
    0xC7CAA854F5D7CDD3UL   /*  682 */,    0x7CAC32883D261D9CUL   /*  683 */,
    0x7690C26423BA942CUL   /*  684 */,    0x17E55524478042B8UL   /*  685 */,
    0xE0BE477656A2389FUL   /*  686 */,    0x4D289B5E67AB2DA0UL   /*  687 */,
    0x44862B9C8FBBFD31UL   /*  688 */,    0xB47CC8049D141365UL   /*  689 */,
    0x822C1B362B91C793UL   /*  690 */,    0x4EB14655FB13DFD8UL   /*  691 */,
    0x1ECBBA0714E2A97BUL   /*  692 */,    0x6143459D5CDE5F14UL   /*  693 */,
    0x53A8FBF1D5F0AC89UL   /*  694 */,    0x97EA04D81C5E5B00UL   /*  695 */,
    0x622181A8D4FDB3F3UL   /*  696 */,    0xE9BCD341572A1208UL   /*  697 */,
    0x1411258643CCE58AUL   /*  698 */,    0x9144C5FEA4C6E0A4UL   /*  699 */,
    0x0D33D06565CF620FUL   /*  700 */,    0x54A48D489F219CA1UL   /*  701 */,
    0xC43E5EAC6D63C821UL   /*  702 */,    0xA9728B3A72770DAFUL   /*  703 */,
    0xD7934E7B20DF87EFUL   /*  704 */,    0xE35503B61A3E86E5UL   /*  705 */,
    0xCAE321FBC819D504UL   /*  706 */,    0x129A50B3AC60BFA6UL   /*  707 */,
    0xCD5E68EA7E9FB6C3UL   /*  708 */,    0xB01C90199483B1C7UL   /*  709 */,
    0x3DE93CD5C295376CUL   /*  710 */,    0xAED52EDF2AB9AD13UL   /*  711 */,
    0x2E60F512C0A07884UL   /*  712 */,    0xBC3D86A3E36210C9UL   /*  713 */,
    0x35269D9B163951CEUL   /*  714 */,    0x0C7D6E2AD0CDB5FAUL   /*  715 */,
    0x59E86297D87F5733UL   /*  716 */,    0x298EF221898DB0E7UL   /*  717 */,
    0x55000029D1A5AA7EUL   /*  718 */,    0x8BC08AE1B5061B45UL   /*  719 */,
    0xC2C31C2B6C92703AUL   /*  720 */,    0x94CC596BAF25EF42UL   /*  721 */,
    0x0A1D73DB22540456UL   /*  722 */,    0x04B6A0F9D9C4179AUL   /*  723 */,
    0xEFFDAFA2AE3D3C60UL   /*  724 */,    0xF7C8075BB49496C4UL   /*  725 */,
    0x9CC5C7141D1CD4E3UL   /*  726 */,    0x78BD1638218E5534UL   /*  727 */,
    0xB2F11568F850246AUL   /*  728 */,    0xEDFABCFA9502BC29UL   /*  729 */,
    0x796CE5F2DA23051BUL   /*  730 */,    0xAAE128B0DC93537CUL   /*  731 */,
    0x3A493DA0EE4B29AEUL   /*  732 */,    0xB5DF6B2C416895D7UL   /*  733 */,
    0xFCABBD25122D7F37UL   /*  734 */,    0x70810B58105DC4B1UL   /*  735 */,
    0xE10FDD37F7882A90UL   /*  736 */,    0x524DCAB5518A3F5CUL   /*  737 */,
    0x3C9E85878451255BUL   /*  738 */,    0x4029828119BD34E2UL   /*  739 */,
    0x74A05B6F5D3CECCBUL   /*  740 */,    0xB610021542E13ECAUL   /*  741 */,
    0x0FF979D12F59E2ACUL   /*  742 */,    0x6037DA27E4F9CC50UL   /*  743 */,
    0x5E92975A0DF1847DUL   /*  744 */,    0xD66DE190D3E623FEUL   /*  745 */,
    0x5032D6B87B568048UL   /*  746 */,    0x9A36B7CE8235216EUL   /*  747 */,
    0x80272A7A24F64B4AUL   /*  748 */,    0x93EFED8B8C6916F7UL   /*  749 */,
    0x37DDBFF44CCE1555UL   /*  750 */,    0x4B95DB5D4B99BD25UL   /*  751 */,
    0x92D3FDA169812FC0UL   /*  752 */,    0xFB1A4A9A90660BB6UL   /*  753 */,
    0x730C196946A4B9B2UL   /*  754 */,    0x81E289AA7F49DA68UL   /*  755 */,
    0x64669A0F83B1A05FUL   /*  756 */,    0x27B3FF7D9644F48BUL   /*  757 */,
    0xCC6B615C8DB675B3UL   /*  758 */,    0x674F20B9BCEBBE95UL   /*  759 */,
    0x6F31238275655982UL   /*  760 */,    0x5AE488713E45CF05UL   /*  761 */,
    0xBF619F9954C21157UL   /*  762 */,    0xEABAC46040A8EAE9UL   /*  763 */,
    0x454C6FE9F2C0C1CDUL   /*  764 */,    0x419CF6496412691CUL   /*  765 */,
    0xD3DC3BEF265B0F70UL   /*  766 */,    0x6D0E60F5C3578A9EUL   /*  767 */,
    0x5B0E608526323C55UL   /*  768 */,    0x1A46C1A9FA1B59F5UL   /*  769 */,
    0xA9E245A17C4C8FFAUL   /*  770 */,    0x65CA5159DB2955D7UL   /*  771 */,
    0x05DB0A76CE35AFC2UL   /*  772 */,    0x81EAC77EA9113D45UL   /*  773 */,
    0x528EF88AB6AC0A0DUL   /*  774 */,    0xA09EA253597BE3FFUL   /*  775 */,
    0x430DDFB3AC48CD56UL   /*  776 */,    0xC4B3A67AF45CE46FUL   /*  777 */,
    0x4ECECFD8FBE2D05EUL   /*  778 */,    0x3EF56F10B39935F0UL   /*  779 */,
    0x0B22D6829CD619C6UL   /*  780 */,    0x17FD460A74DF2069UL   /*  781 */,
    0x6CF8CC8E8510ED40UL   /*  782 */,    0xD6C824BF3A6ECAA7UL   /*  783 */,
    0x61243D581A817049UL   /*  784 */,    0x048BACB6BBC163A2UL   /*  785 */,
    0xD9A38AC27D44CC32UL   /*  786 */,    0x7FDDFF5BAAF410ABUL   /*  787 */,
    0xAD6D495AA804824BUL   /*  788 */,    0xE1A6A74F2D8C9F94UL   /*  789 */,
    0xD4F7851235DEE8E3UL   /*  790 */,    0xFD4B7F886540D893UL   /*  791 */,
    0x247C20042AA4BFDAUL   /*  792 */,    0x096EA1C517D1327CUL   /*  793 */,
    0xD56966B4361A6685UL   /*  794 */,    0x277DA5C31221057DUL   /*  795 */,
    0x94D59893A43ACFF7UL   /*  796 */,    0x64F0C51CCDC02281UL   /*  797 */,
    0x3D33BCC4FF6189DBUL   /*  798 */,    0xE005CB184CE66AF1UL   /*  799 */,
    0xFF5CCD1D1DB99BEAUL   /*  800 */,    0xB0B854A7FE42980FUL   /*  801 */,
    0x7BD46A6A718D4B9FUL   /*  802 */,    0xD10FA8CC22A5FD8CUL   /*  803 */,
    0xD31484952BE4BD31UL   /*  804 */,    0xC7FA975FCB243847UL   /*  805 */,
    0x4886ED1E5846C407UL   /*  806 */,    0x28CDDB791EB70B04UL   /*  807 */,
    0xC2B00BE2F573417FUL   /*  808 */,    0x5C9590452180F877UL   /*  809 */,
    0x7A6BDDFFF370EB00UL   /*  810 */,    0xCE509E38D6D9D6A4UL   /*  811 */,
    0xEBEB0F00647FA702UL   /*  812 */,    0x1DCC06CF76606F06UL   /*  813 */,
    0xE4D9F28BA286FF0AUL   /*  814 */,    0xD85A305DC918C262UL   /*  815 */,
    0x475B1D8732225F54UL   /*  816 */,    0x2D4FB51668CCB5FEUL   /*  817 */,
    0xA679B9D9D72BBA20UL   /*  818 */,    0x53841C0D912D43A5UL   /*  819 */,
    0x3B7EAA48BF12A4E8UL   /*  820 */,    0x781E0E47F22F1DDFUL   /*  821 */,
    0xEFF20CE60AB50973UL   /*  822 */,    0x20D261D19DFFB742UL   /*  823 */,
    0x16A12B03062A2E39UL   /*  824 */,    0x1960EB2239650495UL   /*  825 */,
    0x251C16FED50EB8B8UL   /*  826 */,    0x9AC0C330F826016EUL   /*  827 */,
    0xED152665953E7671UL   /*  828 */,    0x02D63194A6369570UL   /*  829 */,
    0x5074F08394B1C987UL   /*  830 */,    0x70BA598C90B25CE1UL   /*  831 */,
    0x794A15810B9742F6UL   /*  832 */,    0x0D5925E9FCAF8C6CUL   /*  833 */,
    0x3067716CD868744EUL   /*  834 */,    0x910AB077E8D7731BUL   /*  835 */,
    0x6A61BBDB5AC42F61UL   /*  836 */,    0x93513EFBF0851567UL   /*  837 */,
    0xF494724B9E83E9D5UL   /*  838 */,    0xE887E1985C09648DUL   /*  839 */,
    0x34B1D3C675370CFDUL   /*  840 */,    0xDC35E433BC0D255DUL   /*  841 */,
    0xD0AAB84234131BE0UL   /*  842 */,    0x08042A50B48B7EAFUL   /*  843 */,
    0x9997C4EE44A3AB35UL   /*  844 */,    0x829A7B49201799D0UL   /*  845 */,
    0x263B8307B7C54441UL   /*  846 */,    0x752F95F4FD6A6CA6UL   /*  847 */,
    0x927217402C08C6E5UL   /*  848 */,    0x2A8AB754A795D9EEUL   /*  849 */,
    0xA442F7552F72943DUL   /*  850 */,    0x2C31334E19781208UL   /*  851 */,
    0x4FA98D7CEAEE6291UL   /*  852 */,    0x55C3862F665DB309UL   /*  853 */,
    0xBD0610175D53B1F3UL   /*  854 */,    0x46FE6CB840413F27UL   /*  855 */,
    0x3FE03792DF0CFA59UL   /*  856 */,    0xCFE700372EB85E8FUL   /*  857 */,
    0xA7BE29E7ADBCE118UL   /*  858 */,    0xE544EE5CDE8431DDUL   /*  859 */,
    0x8A781B1B41F1873EUL   /*  860 */,    0xA5C94C78A0D2F0E7UL   /*  861 */,
    0x39412E2877B60728UL   /*  862 */,    0xA1265EF3AFC9A62CUL   /*  863 */,
    0xBCC2770C6A2506C5UL   /*  864 */,    0x3AB66DD5DCE1CE12UL   /*  865 */,
    0xE65499D04A675B37UL   /*  866 */,    0x7D8F523481BFD216UL   /*  867 */,
    0x0F6F64FCEC15F389UL   /*  868 */,    0x74EFBE618B5B13C8UL   /*  869 */,
    0xACDC82B714273E1DUL   /*  870 */,    0xDD40BFE003199D17UL   /*  871 */,
    0x37E99257E7E061F8UL   /*  872 */,    0xFA52626904775AAAUL   /*  873 */,
    0x8BBBF63A463D56F9UL   /*  874 */,    0xF0013F1543A26E64UL   /*  875 */,
    0xA8307E9F879EC898UL   /*  876 */,    0xCC4C27A4150177CCUL   /*  877 */,
    0x1B432F2CCA1D3348UL   /*  878 */,    0xDE1D1F8F9F6FA013UL   /*  879 */,
    0x606602A047A7DDD6UL   /*  880 */,    0xD237AB64CC1CB2C7UL   /*  881 */,
    0x9B938E7225FCD1D3UL   /*  882 */,    0xEC4E03708E0FF476UL   /*  883 */,
    0xFEB2FBDA3D03C12DUL   /*  884 */,    0xAE0BCED2EE43889AUL   /*  885 */,
    0x22CB8923EBFB4F43UL   /*  886 */,    0x69360D013CF7396DUL   /*  887 */,
    0x855E3602D2D4E022UL   /*  888 */,    0x073805BAD01F784CUL   /*  889 */,
    0x33E17A133852F546UL   /*  890 */,    0xDF4874058AC7B638UL   /*  891 */,
    0xBA92B29C678AA14AUL   /*  892 */,    0x0CE89FC76CFAADCDUL   /*  893 */,
    0x5F9D4E0908339E34UL   /*  894 */,    0xF1AFE9291F5923B9UL   /*  895 */,
    0x6E3480F60F4A265FUL   /*  896 */,    0xEEBF3A2AB29B841CUL   /*  897 */,
    0xE21938A88F91B4ADUL   /*  898 */,    0x57DFEFF845C6D3C3UL   /*  899 */,
    0x2F006B0BF62CAAF2UL   /*  900 */,    0x62F479EF6F75EE78UL   /*  901 */,
    0x11A55AD41C8916A9UL   /*  902 */,    0xF229D29084FED453UL   /*  903 */,
    0x42F1C27B16B000E6UL   /*  904 */,    0x2B1F76749823C074UL   /*  905 */,
    0x4B76ECA3C2745360UL   /*  906 */,    0x8C98F463B91691BDUL   /*  907 */,
    0x14BCC93CF1ADE66AUL   /*  908 */,    0x8885213E6D458397UL   /*  909 */,
    0x8E177DF0274D4711UL   /*  910 */,    0xB49B73B5503F2951UL   /*  911 */,
    0x10168168C3F96B6BUL   /*  912 */,    0x0E3D963B63CAB0AEUL   /*  913 */,
    0x8DFC4B5655A1DB14UL   /*  914 */,    0xF789F1356E14DE5CUL   /*  915 */,
    0x683E68AF4E51DAC1UL   /*  916 */,    0xC9A84F9D8D4B0FD9UL   /*  917 */,
    0x3691E03F52A0F9D1UL   /*  918 */,    0x5ED86E46E1878E80UL   /*  919 */,
    0x3C711A0E99D07150UL   /*  920 */,    0x5A0865B20C4E9310UL   /*  921 */,
    0x56FBFC1FE4F0682EUL   /*  922 */,    0xEA8D5DE3105EDF9BUL   /*  923 */,
    0x71ABFDB12379187AUL   /*  924 */,    0x2EB99DE1BEE77B9CUL   /*  925 */,
    0x21ECC0EA33CF4523UL   /*  926 */,    0x59A4D7521805C7A1UL   /*  927 */,
    0x3896F5EB56AE7C72UL   /*  928 */,    0xAA638F3DB18F75DCUL   /*  929 */,
    0x9F39358DABE9808EUL   /*  930 */,    0xB7DEFA91C00B72ACUL   /*  931 */,
    0x6B5541FD62492D92UL   /*  932 */,    0x6DC6DEE8F92E4D5BUL   /*  933 */,
    0x353F57ABC4BEEA7EUL   /*  934 */,    0x735769D6DA5690CEUL   /*  935 */,
    0x0A234AA642391484UL   /*  936 */,    0xF6F9508028F80D9DUL   /*  937 */,
    0xB8E319A27AB3F215UL   /*  938 */,    0x31AD9C1151341A4DUL   /*  939 */,
    0x773C22A57BEF5805UL   /*  940 */,    0x45C7561A07968633UL   /*  941 */,
    0xF913DA9E249DBE36UL   /*  942 */,    0xDA652D9B78A64C68UL   /*  943 */,
    0x4C27A97F3BC334EFUL   /*  944 */,    0x76621220E66B17F4UL   /*  945 */,
    0x967743899ACD7D0BUL   /*  946 */,    0xF3EE5BCAE0ED6782UL   /*  947 */,
    0x409F753600C879FCUL   /*  948 */,    0x06D09A39B5926DB6UL   /*  949 */,
    0x6F83AEB0317AC588UL   /*  950 */,    0x01E6CA4A86381F21UL   /*  951 */,
    0x66FF3462D19F3025UL   /*  952 */,    0x72207C24DDFD3BFBUL   /*  953 */,
    0x4AF6B6D3E2ECE2EBUL   /*  954 */,    0x9C994DBEC7EA08DEUL   /*  955 */,
    0x49ACE597B09A8BC4UL   /*  956 */,    0xB38C4766CF0797BAUL   /*  957 */,
    0x131B9373C57C2A75UL   /*  958 */,    0xB1822CCE61931E58UL   /*  959 */,
    0x9D7555B909BA1C0CUL   /*  960 */,    0x127FAFDD937D11D2UL   /*  961 */,
    0x29DA3BADC66D92E4UL   /*  962 */,    0xA2C1D57154C2ECBCUL   /*  963 */,
    0x58C5134D82F6FE24UL   /*  964 */,    0x1C3AE3515B62274FUL   /*  965 */,
    0xE907C82E01CB8126UL   /*  966 */,    0xF8ED091913E37FCBUL   /*  967 */,
    0x3249D8F9C80046C9UL   /*  968 */,    0x80CF9BEDE388FB63UL   /*  969 */,
    0x1881539A116CF19EUL   /*  970 */,    0x5103F3F76BD52457UL   /*  971 */,
    0x15B7E6F5AE47F7A8UL   /*  972 */,    0xDBD7C6DED47E9CCFUL   /*  973 */,
    0x44E55C410228BB1AUL   /*  974 */,    0xB647D4255EDB4E99UL   /*  975 */,
    0x5D11882BB8AAFC30UL   /*  976 */,    0xF5098BBB29D3212AUL   /*  977 */,
    0x8FB5EA14E90296B3UL   /*  978 */,    0x677B942157DD025AUL   /*  979 */,
    0xFB58E7C0A390ACB5UL   /*  980 */,    0x89D3674C83BD4A01UL   /*  981 */,
    0x9E2DA4DF4BF3B93BUL   /*  982 */,    0xFCC41E328CAB4829UL   /*  983 */,
    0x03F38C96BA582C52UL   /*  984 */,    0xCAD1BDBD7FD85DB2UL   /*  985 */,
    0xBBB442C16082AE83UL   /*  986 */,    0xB95FE86BA5DA9AB0UL   /*  987 */,
    0xB22E04673771A93FUL   /*  988 */,    0x845358C9493152D8UL   /*  989 */,
    0xBE2A488697B4541EUL   /*  990 */,    0x95A2DC2DD38E6966UL   /*  991 */,
    0xC02C11AC923C852BUL   /*  992 */,    0x2388B1990DF2A87BUL   /*  993 */,
    0x7C8008FA1B4F37BEUL   /*  994 */,    0x1F70D0C84D54E503UL   /*  995 */,
    0x5490ADEC7ECE57D4UL   /*  996 */,    0x002B3C27D9063A3AUL   /*  997 */,
    0x7EAEA3848030A2BFUL   /*  998 */,    0xC602326DED2003C0UL   /*  999 */,
    0x83A7287D69A94086UL   /* 1000 */,    0xC57A5FCB30F57A8AUL   /* 1001 */,
    0xB56844E479EBE779UL   /* 1002 */,    0xA373B40F05DCBCE9UL   /* 1003 */,
    0xD71A786E88570EE2UL   /* 1004 */,    0x879CBACDBDE8F6A0UL   /* 1005 */,
    0x976AD1BCC164A32FUL   /* 1006 */,    0xAB21E25E9666D78BUL   /* 1007 */,
    0x901063AAE5E5C33CUL   /* 1008 */,    0x9818B34448698D90UL   /* 1009 */,
    0xE36487AE3E1E8ABBUL   /* 1010 */,    0xAFBDF931893BDCB4UL   /* 1011 */,
    0x6345A0DC5FBBD519UL   /* 1012 */,    0x8628FE269B9465CAUL   /* 1013 */,
    0x1E5D01603F9C51ECUL   /* 1014 */,    0x4DE44006A15049B7UL   /* 1015 */,
    0xBF6C70E5F776CBB1UL   /* 1016 */,    0x411218F2EF552BEDUL   /* 1017 */,
    0xCB0C0708705A36A3UL   /* 1018 */,    0xE74D14754F986044UL   /* 1019 */,
    0xCD56D9430EA8280EUL   /* 1020 */,    0xC12591D7535F5065UL   /* 1021 */,
    0xC83223F1720AEF96UL   /* 1022 */,    0xC3A0396F7363A51FUL   /* 1023 */};

#define TTBLOCK_SIZE 1024

#define TTBLOCK_SIZE_P1 (TTBLOCK_SIZE+1)
#define LTH_BYTES 25
#define ITH_BYTES 49

#define POW2(exp) (1 << exp)

#define memcpy_macro(__adrr_spc1, __adrr_spc2) \
  uint i; \
  for (i = 0; i < sz; ++i) *((__adrr_spc1 uchar *)dst + i) = *((__adrr_spc2 uchar *)src + i);

//void memcpy_pp(void *dst, void *src, size_t sz) { memcpy_macro(__private, __private) }
void memcpy_pg(void *dst, __global void *src, size_t sz) { memcpy_macro(__private, __global) }
void memcpy_gp(__global void *dst, void *src, size_t sz) { memcpy_macro(__global, __private) }
void memcpy_gg(__global void *dst, __global void *src, size_t sz) { memcpy_macro(__global, __global) }

/*
 * Copy block from src to dst.
 * dst size must be >= blk_sz/sizeof(word64)
 * src size must be >= blk_sz/sizeof(word64)
 * blk_sz must be multiple of sizeof(word64)
 */
void blkcpy_gg(__global word64 *dst, __global word64 *src, word64 blk_sz)
{
  word64 i;
  for (i = 0; i < blk_sz/sizeof(word64); ++i) dst[i] = src[i];
}

/*
 * Copy block from src to (byte *)dst + 1.
 * dst size must be >= blk_sz/sizeof(word64)+1
 * src size must be >= blk_sz/sizeof(word64)
 * blk_sz must be multiple of sizeof(word64)
 */
void blkcpy_lp_gg(__global word64 *dst, __global word64 *src, word64 blk_sz)
{
  word64 i;

#ifdef LITTLE_ENDIAN
  dst[0] = (src[0] & 0x00ffffffffffffff)<<8 | dst[0]&0x00000000000000ff;
#else
  dst[0] = (src[0] & 0xffffffffffffff00)>>8 | dst[0]&0xff00000000000000;
#endif

  for (i = 1; i < blk_sz/sizeof(word64); ++i) {
#ifdef LITTLE_ENDIAN
    dst[i] = (src[i]&0x00ffffffffffffff)<<8 | (src[i-1]&0xff00000000000000)>>56;
#else
    dst[i] = (src[i]&0xffffffffffffff00)>>8 | (src[i-1]&0x00000000000000ff)<<56;
#endif
  }

#ifdef LITTLE_ENDIAN
  dst[i] = (src[i-1] & 0xff00000000000000)>>56 | dst[i]&0xffffffffffffff00;
#else
  dst[i] = (src[i-1] & 0x00000000000000ff)<<56 | dst[i]&0x00ffffffffffffff;
#endif
}

void blkcpy_lp_pg(word64 *dst, __global word64 *src, word64 blk_sz)
{
  word64 i;

#ifdef LITTLE_ENDIAN
  dst[0] = (src[0] & 0x00ffffffffffffff)<<8 | dst[0]&0x00000000000000ff;
#else
  dst[0] = (src[0] & 0xffffffffffffff00)>>8 | dst[0]&0xff00000000000000;
#endif

  for (i = 1; i < blk_sz/sizeof(word64); ++i) {
#ifdef LITTLE_ENDIAN
    dst[i] = (src[i]&0x00ffffffffffffff)<<8 | (src[i-1]&0xff00000000000000)>>56;
#else
    dst[i] = (src[i]&0xffffffffffffff00)>>8 | (src[i-1]&0x00000000000000ff)<<56;
#endif
  }

#ifdef LITTLE_ENDIAN
  dst[i] = (src[i-1] & 0xff00000000000000)>>56 | dst[i]&0xffffffffffffff00;
#else
  dst[i] = (src[i-1] & 0x00000000000000ff)<<56 | dst[i]&0x00ffffffffffffff;
#endif
}

void blkcpy_lp_gp(__global word64 *dst, word64 *src, word64 blk_sz)
{
  word64 i;

#ifdef LITTLE_ENDIAN
  dst[0] = (src[0] & 0x00ffffffffffffff)<<8 | dst[0]&0x00000000000000ff;
#else
  dst[0] = (src[0] & 0xffffffffffffff00)>>8 | dst[0]&0xff00000000000000;
#endif

  for (i = 1; i < blk_sz/sizeof(word64); ++i) {
#ifdef LITTLE_ENDIAN
    dst[i] = (src[i]&0x00ffffffffffffff)<<8 | (src[i-1]&0xff00000000000000)>>56;
#else
    dst[i] = (src[i]&0xffffffffffffff00)>>8 | (src[i-1]&0x00000000000000ff)<<56;
#endif
  }

#ifdef LITTLE_ENDIAN
  dst[i] = (src[i-1] & 0xff00000000000000)>>56 | dst[i]&0xffffffffffffff00;
#else
  dst[i] = (src[i-1] & 0x00000000000000ff)<<56 | dst[i]&0x00ffffffffffffff;
#endif
}

/*
 * Computes Tiger Tree Root hash.
 * __in_out data - input data buffer
 * Assumptions:
 * 1. data buffer size is multiple of TTBLOCK_SIZE
 * even last_blk_sz < TTBLOCK_SIZE;
 * 2. TTBLOCK_SIZE >= 2*TH_BYTES + 1);
 * __in blk_cnt - blocks count;
 * __in last_blk_sz - size of the last input block (in bytes)
 * Resulting root hash is placed in memory pointed by res.
 */
__kernel void
ttr(__global uchar *data, ulong blk_cnt, ulong last_blk_sz, __global word64 res[3])
{
  ulong i, l;
  uchar flg;
  word64 tmp[TTBLOCK_SIZE/sizeof(word64) + 1];
  word64 th[TH_BYTES/sizeof(word64)]; // uses when needed word64-aligned buffer
  
  size_t gid, gsz;
  
  gid = get_global_id(0);
  gsz = get_global_size(0);

  tmp[0] = 0;
  l = 0;

  for (i = 2*gid + 1; i < blk_cnt-1; i += 2*gsz) {
    blkcpy_lp_pg(tmp, data + TTBLOCK_SIZE*(i-1), TTBLOCK_SIZE);
	tiger_pg(tmp, TTBLOCK_SIZE_P1, data+TTBLOCK_SIZE*(i-1));

    blkcpy_lp_pg(tmp, data + TTBLOCK_SIZE*i, TTBLOCK_SIZE);
	data[TTBLOCK_SIZE*i] = 1;
	tiger_pp(tmp, TTBLOCK_SIZE_P1, th);
	blkcpy_lp_gp(data+TTBLOCK_SIZE*i+TH_BYTES, th, TH_BYTES);
  }
  
  if (gid == gsz-1) {
    if (blk_cnt & 1) {
      i = blk_cnt - 1;

      memcpy_pg((byte *)tmp + 1, data + TTBLOCK_SIZE*i, last_blk_sz);
      tiger_pg(tmp, last_blk_sz+1, data + TTBLOCK_SIZE*i);
    } else {
      i = blk_cnt - 2;

      blkcpy_lp_pg(tmp, data + TTBLOCK_SIZE*i, TTBLOCK_SIZE);
      tiger_pg(tmp, TTBLOCK_SIZE_P1, data + TTBLOCK_SIZE*i);

      ++i;
	
      memcpy_pg((byte *)tmp + 1, data + TTBLOCK_SIZE*i, last_blk_sz);
      tiger_pp(tmp, last_blk_sz+1, th);
	  blkcpy_lp_gp(data+TTBLOCK_SIZE*i+TH_BYTES, th, TH_BYTES);
	
      data[TTBLOCK_SIZE*i] = 1;
    }
  }

  barrier(CLK_GLOBAL_MEM_FENCE);

  /* Here all bottom blocks are hashed */
  
  while (blk_cnt > 1) {
    ++l;
    flg = blk_cnt & 1;
    blk_cnt = (blk_cnt+1) / 2;

    for (i = 2*gid + 1; i < blk_cnt-1; i += 2*gsz) {
	  blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
        data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)),
		TH_BYTES);
		
	  blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)+1),
        data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)),
		TH_BYTES);
		
	  tiger_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
	    ITH_BYTES,
	    data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1))
	  );
		
	  tiger_gp(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)+1),
	    ITH_BYTES,
	    th
      );
	  blkcpy_lp_gp(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i))+TH_BYTES, th, TH_BYTES);
	  
	  data[TTBLOCK_SIZE*POW2(l-1)*(2*(i))] = 1;
	}
	
    if (gid == gsz-1) {
	  if ( blk_cnt & 1 ) {
        i = blk_cnt;

	    if ( !flg ) {
	      blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
            data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)),
		    TH_BYTES);
		
	      tiger_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
    	    ITH_BYTES,
	        data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1))
	      );
	    }
      } else {
        i = blk_cnt - 1;

	    if ( flg ) {
	      blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
            data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)),
		    TH_BYTES);
		
	      blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)) + TH_BYTES,
            data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)),
	        TH_BYTES);
		
	      tiger_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
	        ITH_BYTES,
	        data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1))
	      );
	  
	      data[TTBLOCK_SIZE*POW2(l-1)*(2*(i))] = 1;
	    } else {
	      blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
            data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)),
		    TH_BYTES);
		
	      blkcpy_lp_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)+1),
            data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)),
	        TH_BYTES);
		
	      tiger_gg(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1)+1),
	        ITH_BYTES,
	        data + TTBLOCK_SIZE*POW2(l-1)*(2*(i-1))
	      );
		
	      tiger_gp(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i)+1),
	        ITH_BYTES,
	        th
          );
	      blkcpy_lp_gp(data + TTBLOCK_SIZE*POW2(l-1)*(2*(i))+TH_BYTES, th, TH_BYTES);
	  
	      data[TTBLOCK_SIZE*POW2(l-1)*(2*(i))] = 1;
	    }
	  }
	}
	
    barrier(CLK_GLOBAL_MEM_FENCE);
  }
  
  if (gid == gsz-1) blkcpy_gg(res, data, TH_BYTES);
}